
#include "media_gpu.h"
#include <assert.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <helper_cuda.h>
#include <device_functions.h> //_synchronizedThreads()
#include <stdlib.h>
#include <stdio.h>
#include <time.h>
#include <helper_image.h>
#include <sys/time.h>
#define SHARED_MEMX 64
#define SHARED_MEMY 64
extern __shared__ short unsigned int sumX[];

/**
 * add the kx rows and store the result in sumX
 * @param img
 * @param imgX image size in X
 * @param kx window size in X
 * @param ky window size in Y
 * @param dimX
 * @param dimY
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 *///7+8*(4+11+11*4+12+6) = 623
inline __device__ void sumRow(uchar4* img, int imgX, int kx,int ky,int dimX,int dimY, int pixelsPerThread){
	int kxMed = (int)kx/2;//1
	int blockY = (blockIdx.y*dimY);//1
	int blockX = (blockIdx.x*dimX);//1
	int length = dimX*dimY;//1
	bool isFirst=true;
	uchar4 first;
	uchar4 p;

	if(blockIdx.y!=0)
			blockY = (blockIdx.y*dimY)-blockIdx.y*ky;//3
	for(int X=threadIdx.x*pixelsPerThread;X<min((threadIdx.x+1)*pixelsPerThread,imgX);X++){//8*4*(
		int idx =(blockY+threadIdx.y)*imgX+blockX+X;//5
		int sIdx = threadIdx.y*dimX+X-1;//3
		int sum1=0,sum2=0,sum3=0;
		if(blockX+X>=kxMed && blockX+X<imgX-kxMed){//3
			if(isFirst){
				for(int i=-kxMed;i<=kxMed;i++){//11*(
					p = img[idx+i];//1
					sum1 += p.x;//1
					sum2 += p.y;//1
					sum3 += p.z;//1
				}
				isFirst=false;
			}
			else{
				p = img[idx+kxMed];//1
				first = img[idx-kxMed-1];//2

				sum1 = sumX[sIdx]-first.x+p.x;//2
				sum2 = sumX[length+sIdx]-first.y+p.y;//3
				sum3 = sumX[2*length+sIdx]-first.z+p.z;//4
			}
			sumX[sIdx+1] = sum1;//1
			sumX[length+sIdx+1] = sum2;//2
			sumX[2*length+sIdx+1] = sum3;//3
		}
	}
}
/**
 * add ky columns and store the result in buf
 * @param buf image
 * @param imgX image size in X
 * @param imgY image size in Y
 * @param ky window size in Y
 * @param kx window size in X
 * @param dimX
 * @param dimY
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 *///9+8*(4+6+11*9+8)=945
inline __device__ void sumColumn(uchar4* buf, int imgX, int imgY, int ky,int kx,int dimX,int dimY, int pixelsPerThread){
	int t = (ky*kx);//1
	int kyMed = (int)ky/2;//1
	int kxMed = (int)kx/2;//1
	int blockY = (blockIdx.y*dimY);//1
	int blockX = (blockIdx.x*dimX);//1
	int length = dimX*dimY;//1
	if(blockIdx.y!=0)
		blockY = (blockIdx.y*dimY)-blockIdx.y*ky;//3
	for(int X=threadIdx.x*pixelsPerThread ;X<min((threadIdx.x+1)*pixelsPerThread,imgX);X++){//8*4*(
		int sum1=0,sum2=0,sum3=0;//6
		if(blockX+X>=kxMed && blockX+X<imgX-kxMed && threadIdx.y>=kyMed && threadIdx.y<dimY-kyMed && blockY+threadIdx.y<imgY-kyMed){
			uchar4 p;
			for(int j=-kyMed;j<=kyMed;j++){//11*(
				int sIdx=(threadIdx.y+j)*dimX+X;//3
				sum1 += sumX[sIdx];//1
				sum2 += sumX[length+sIdx];//2
				sum3 += sumX[2*length+sIdx];//3
			}
			p.x=sum1/t;//1
			p.y=sum2/t;//1
			p.z=sum3/t;//1
			buf[(blockY+threadIdx.y)*imgX+blockX+X] = p;//5
		}
	}
}
/**
 * media filter in cuda
 * @param buf
 * @param img
 * @param imgX image size in X
 * @param imgY image size in Y
 * @param kx window size in X
 * @param ky window size in Y
 * @param dimX image size in X in the block
 * @param dimY image size in Y in the block
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 */
__global__ void sumCentre(uchar4* buf,uchar4* img, int imgX, int imgY, int kx,int ky, int dimX, int dimY, int pixelsPerThread){

	sumRow(img,imgX,kx,ky,dimX,dimY,pixelsPerThread);
	__syncthreads();
	sumColumn(buf, imgX, imgY, ky,kx,dimX,dimY, pixelsPerThread);
}
/**
 * media filter
 * @param h_img image
 * @param width image size in X
 * @param height image size in Y
 * @param dims
 * @param kx window size in X
 * @param ky window size in Y
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 * @return media filter image h_img
 */
uchar4* media_gpu(uchar4* h_img,int width, int height,int* dims, int kx, int ky, int pixelsPerThread){

	assert(h_img!=NULL && width!=0 && height!=0);

	uchar4* result = (uchar4*) calloc(sizeof(uchar4),width*height);
	//DEVICE
	uchar4* d_img;
	uchar4* d_buf;
	//ALOCAR MEMORIA HOST

	checkCudaErrors(cudaMalloc((void**) &(d_img), width*height*sizeof(uchar4)));
	checkCudaErrors(cudaMalloc((void**) &(d_buf), width*height*sizeof(uchar4)));

	//COPY IMG TO DEVICE
	checkCudaErrors(cudaMemcpy(d_img, h_img, width*height*sizeof(uchar4), cudaMemcpyHostToDevice));
	checkCudaErrors(cudaMemcpy(d_buf, d_img, width*height*sizeof(uchar4), cudaMemcpyDeviceToDevice));

	int mem = 3*SHARED_MEMX*SHARED_MEMY*sizeof(short unsigned int);

	dim3 dimGrid(width/(SHARED_MEMY-kx-1)+1,height/(SHARED_MEMY-ky-1)+1,1);
	dim3 threads(SHARED_MEMX/pixelsPerThread,min(SHARED_MEMY,height),1);
	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start);
	sumCentre<<<dimGrid,threads,mem>>>(d_buf,d_img,width,height, kx,ky, SHARED_MEMX,SHARED_MEMY,pixelsPerThread);
	cudaEventRecord(stop);
	cudaEventSynchronize(stop);
	checkCudaErrors(cudaMemcpy(result, d_buf, width*height*sizeof(uchar4), cudaMemcpyDeviceToHost));


	float milliseconds = 0;
	cudaEventElapsedTime(&milliseconds, start, stop);
	long int hilos = dimGrid.x*dimGrid.y*threads.x*threads.y;
	printf("GFLOPS/s: %lf\n", (hilos*1568)/(milliseconds*pow(10,6)));
	printf("Effective Bandwidth (GB/s): %f\n", hilos*pixelsPerThread*kx*sizeof(uchar4)/(milliseconds*pow(10,6)));

	checkCudaErrors(cudaFree(d_img));
	checkCudaErrors(cudaFree(d_buf));
	return result;
}
